#include <iostream>
#include <cuda_runtime.h>
#include "../common/common.h"
#define DEVICENUM 0

__global__ void norm_kernel(float *inputValue, float *res)
{
    int idx = threadIdx.x + blockDim.x * blockIdx.x;
    int temp = *inputValue;
    temp = temp + 1;
    *inputValue = temp;
    res[idx] = *inputValue;
}

__global__ void atomic_kernel(float *inputValue, float *res)
{
    int idx = threadIdx.x + blockDim.x * blockIdx.x;
    // 如果启动32个线程，存储在*inputValue所指位置中的值应该是32
    int temp = atomicAdd(inputValue, 1);
    res[idx] = *inputValue;
}

int main(int argc, char **argv)
{
    cudaSetDevice(DEVICENUM);

    int nElem = 32;
    size_t nBytes = nElem * sizeof(float);

    float *host_array_res = (float *)malloc(nBytes);
    memset(host_array_res, 0, nBytes);

    float *input_value = nullptr;
    float *device_array_out = nullptr;
    CHECK(cudaMalloc((float **)&input_value, sizeof(float)));
    CHECK(cudaMalloc((float **)&device_array_out, nBytes));
    CHECK(cudaMemset(input_value, 0.0f, sizeof(float)));
    CHECK(cudaMemset(device_array_out, 0.0f, nBytes));

    norm_kernel<<<1, nElem>>>(input_value, device_array_out);
    CHECK(cudaDeviceSynchronize());
    CHECK(cudaMemcpy(host_array_res, device_array_out, nBytes, cudaMemcpyDeviceToHost));
    printVector(host_array_res, nElem, "host_array_res");

    CHECK(cudaMemset(input_value, 0.0f, sizeof(float)));

    atomic_kernel<<<1, nElem>>>(input_value, device_array_out);
    CHECK(cudaDeviceSynchronize());
    CHECK(cudaMemcpy(host_array_res, device_array_out, nBytes, cudaMemcpyDeviceToHost));
    printVector(host_array_res, nElem, "host_array_res");

    free(host_array_res);
    cudaFree(input_value);
    cudaFree(device_array_out);

    return 0;
}